S51205 From the Macro to the Micro: CUDA Developer Tools Find and Fix Problems at Any Scale

Jackson Marusarz | Technical Product Manager
March 22nd 2023

目录 (Table of Contents)

议程 (Agenda)

本演示将介绍以下内容：

议程概览
议程概览

上图展示了从宏观（服务器机架）到微观（GPU）的层级，强调了不同层面的工具应用。

开发者工具生态系统 (Developer Tools Ecosystem)

NVIDIA 提供了一整套开发者工具，涵盖了调试器、性能分析器、正确性检查器和IDE集成。

调试器 (Debuggers)

性能分析器 (Profilers)

正确性检查器 (Correctness Checker)

IDE 集成 (IDE Integrations)

开发者工具生态系统概览
开发者工具生态系统概览

Grace 及其他 Arm 平台上的 GUI 支持 (GUI Support for Grace and other Arm Platforms)

NVIDIA 的调试和分析工具现在支持在 Arm 平台上原生运行 GUI。
* Nsight Systems - Nsight Compute - Nsight Visual Studio Code Edition
* GUI 可原生运行在 NVIDIA® Jetson AGX Orin™ SoC 和 Arm 服务器平台上。
* 可以使用新的原生 GUI 或现有的远程收集功能。

Arm平台上的GUI支持
Arm平台上的GUI支持

更新的网站/文档/图标 (Updated Websites/Docs/Icons)

NVIDIA 开发者工具 (NVIDIA Developer Tools) 提供了一套强大的库、SDK和开发工具，适用于桌面和移动目标，帮助开发者利用最先进的加速计算硬件进行构建、调试、分析和开发。

您可以根据工作细分领域（如 CUDA/Compute, Graphics, OptiX, Deep Learning）和开发活动（如 Code Development, Debugging/Correctness, Profiling, Platform Analysis, Kernel Analysis, APIs）来查找相关工具。

Nsight Compute 是一款用于 CUDA 应用的交互式内核性能分析器，提供详细的性能指标和API。

相关网站和文档也已更新，并提供了新的工具图标。

更多信息请访问：https://developer.nvidia.com/tools-overview

更新的网站、文档和图标
更新的网站、文档和图标

计算调试器 (Compute Debuggers)

调试设备上运行的 GPU 内核 (Debug GPU Kernels Running on Device)

计算调试器功能
计算调试器功能

CUDA GDB (Command-line and IDE Back-end Debugger)

CUDA GDB 是一款命令行和 IDE 后端调试器，提供以下功能：
* 统一的 CPU 和 CUDA 调试体验。
* 支持 CUDA-C/SASS。
* 基于 GDB 构建，并使用许多相同的 CLI 命令。
* 支持本地/远程连接。
* 可作为 IDE 调试器的后端。

CUDA GDB 详细信息
CUDA GDB 详细信息

Compute Sanitizer (自动扫描错误和内存问题)

Compute Sanitizer 通过子工具检查正确性问题：
* Memcheck：内存访问错误和内存泄漏检测工具。
* Racecheck：共享内存数据访问危害检测工具。
* Initcheck：未初始化设备全局内存访问检测工具。
* Synccheck：线程同步危害检测工具。

以下是一个 Compute Sanitizer 运行示例，展示了检测到的内存错误和访问越界问题：

$ make run_memcheck
/usr/local/cuda/compute-sanitizer/compute-sanitizer --destroy-on-device-error memcheck_demo
========= COMPUTE-SANITI泽
Mallocing memory
Running unaligned_kernel
Ran unaligned_kernel: no error
Sync: no error
Running out_of_bounds_kernel
Ran out_of_bounds_kernel: no error
Sync: no error
========= Invalid __global__ write of size 4 bytes
    at 0x7e71ac080000 is unaligned
    at 0x6e80 in memcheck_demo.cu:6:unaligned_kernel(void)
    by thread (0,0,0) in block (0,0,0)
    Address 0x400100001 is misaligned

更多示例可在 GitHub 上找到：https://github.com/NVIDIA/compute-sanitizer-samples

Compute Sanitizer 概览
Compute Sanitizer 概览

正确性工具功能 (Correctness Tools Features)

使用统一后端进行调试 (Debuggers using Unified Backend)

新的 Compute Sanitizer 用例 (New Compute Sanitizer use cases)

Nsight Visual Studio Code Edition (Nsight Visual Studio Code Edition)

正确性工具新功能
正确性工具新功能

NVIDIA Tools eXtension (NVTX v3)

NVTX v3 是一种工具扩展 API，旨在通过注解（标记、范围、嵌套范围）来修饰应用程序源代码，以帮助可视化执行过程，从而进行调试、追踪和性能分析。

nvtxMark("This is a marker");
nvtxRangePush("This is a push/pop range");
// do something interesting in the range
nvtxRangePop(); // Pop must be on same thread as corresponding Push
nvtxRangeHandle_t handle = nvtxRangeStart("This is a start/end range");
// Somewhere else in the code, not necessarily same thread as Start call:
nvtxRangeEnd(handle);

API 参考：
* https://nvidia.github.io/NVTX/doxygen/index.html
* https://nvidia.github.io/NVTX/doxygen-cpp/index.html

NVTX v3 功能和用法
NVTX v3 功能和用法

NVIDIA SDKs 和 NVTX - 完整的生态系统 (NVIDIA SDKs and NVTX - A Complete Ecosystem)

NVTX 与各种 NVIDIA SDK 和库集成，形成了一个完整的生态系统，支持全面的性能分析和调试。

NVIDIA SDKs和NVTX生态系统
NVIDIA SDKs和NVTX生态系统

Nsight Systems

主要特性

Nsight Systems 是一款系统性能分析工具，具备以下关键特性：

Nsight Systems 主要特性和概览
在Page 16的右侧，展示了Nsight Systems时间轴视图的三个示例，用以说明其可视化能力。

Nsight Systems 界面概览

Nsight Systems 的可视化界面能够详细展示应用程序的性能数据，包括：
* 进程和线程：显示系统中运行的进程及其线程活动。
* 线程状态：展示各线程的运行状态。
* cuDNN和cuBLAS跟踪：对深度学习库的API调用进行跟踪。
* 内核和内存传输活动：监控GPU内核的执行和内存传输操作。
* 多GPU支持：支持对多GPU系统的性能分析。

Nsight Systems 2023.2 界面概览
Page 17 展示了 Nsight Systems 2023.2 的详细时间轴视图，并用标注指明了不同区域显示的信息类型。

兴趣区域的缩放/过滤

Nsight Systems 提供了强大的缩放和过滤功能，使用户能够精确地聚焦于感兴趣的特定区域。

缩放/过滤至精确兴趣区域
Page 18 展示了如何通过选择区域并执行“Zoom into Selection”操作来放大时间轴上的特定性能事件。

Nsight Systems 新功能

CPU 和 OS 事件

Nsight Systems 扩展了对 CPU 和操作系统事件的分析能力：
* CPU X86_64/SBSA 硬件计数器：
* 包括缓存未命中、DRAM 访问等。

CPU X86_64/SBSA 硬件计数器
Page 20 展示了时间轴视图中对 CPU Cache (如 LL Cache Load Misses, CPU Cycles, Cache Misses) 和 Uncore (如 nvidia_scf_pmu_0/mem_rd_access/, nvidia_scf_pmu_0/mem_rd_data/) 性能指标的详细监控。

Linux OS 事件计数器
Page 21 展示了时间轴视图中对操作系统事件 (如 OS Context Switches, Page Faults) 的详细监控。

Python 调用栈采样

Python 调用栈采样
Page 22 展示了时间轴视图中的 Python Backtrace，以及在下方窗格中显示的详细调用栈信息。

NVIDIA 网络 NIC/DPU 指标采样

NVIDIA 网络 NIC/DPU 指标采样
Page 23 展示了 NIC 0 到 NIC 3 的网络指标，包括发送和接收字节数以及发送等待时间。

网络 API 跟踪 - 扩展描述/工具提示

网络 API 跟踪 - 扩展描述/工具提示
Page 24 展示了“Events View”中 MPI_Bcast 调用的详细信息，包括开始时间 0,164935s，结束时间 0,165045s (+109,210 µs)，线程 ID 1342434，发送字节数 0，接收字节数 4，根节点 0，以及 MPI_COMM_WORLD。

Nsight Systems 多节点分析

Nsight Systems 多节点分析
Page 25 展示了 Nsight Systems 多节点分析的命令行帮助信息 (nsys recipe -help) 和一个 Jupyter Notebook 中展示的分析报告。

Nsight Systems 多节点分析（GPU 利用率地图）
Page 26 展示了使用 cuda_gpu_time_util_map 配方进行多节点分析的结果，通过热力图形式对比了优化“Before”和“After”的 GPU 利用率随时间 (Duration (s)) 的变化。

Nsight Compute

Nsight Compute - 内核分析器

主要特性：

Nsight Compute 主要特性和概览
Page 28 展示了 Nsight Compute 界面，包括 GPU Speed of Light 和 GPU Utilization 的详细指标以及各项性能数据列表。

Nsight Compute GUI 界面

Nsight Compute 的 GUI 界面提供：
* 目标度量指标部分：集中展示关键性能指标。
* 可定制的数据收集和呈现：用户可以根据需求调整数据收集和显示方式。
* 内置专业知识用于指导分析和优化：提供建议和瓶颈识别。

Nsight Compute GUI 界面
Page 29 展示了 Nsight Compute GUI 界面的关键部分，包括“GPU Speed of Light”的性能概览、GPU 利用率图表以及根据内置专业知识提供的“Bottleneck”和“Compute Workload Analysis”建议。

内存工作负载分析

Nsight Compute 提供深入的内存工作负载分析，包括：
* 可视化内存分析图表。
* 峰值性能比率的度量指标。

内存工作负载分析
Page 30 展示了一个详细的内存分析图表，可视化了内存层次结构（Global, Local, Texture, Surface, Shared）及其与缓存（Unified Cache, L2 Cache）和主存（System Memory, Device Memory）之间的交互。图表下方还列出了共享内存和一级缓存的详细性能指标。

Nsight Compute 详细分析

源代码/PTX/SASS 分析与关联

Nsight Compute 提供了深入的源代码、PTX（并行线程执行）和 SASS（汇编）层面的性能分析和关联。它通过度量指标热力图快速识别性能热点，并为每条指令提供详细的度量指标。

Nsight Compute Source/PTX/SASS分析与关联
Nsight Compute Source/PTX/SASS分析与关联

图示信息：
* 总采样计数: 111
* 屏障: 43 (38.7%)
* Mio 节流: 21 (18.9%)
* 未选择: 8 (7.2%)
* 已选择: 7 (6.3%)
* 短记分牌: 16 (14.4%)
* 等待: 16 (14.4%)

Nsight Compute 新功能

Nsight Systems 的集成基本跟踪

Nsight Compute 与 Nsight Systems 集成，提供了一个从系统级别到内核级别的无缝分析流程。
* 在连接对话框中使用“System Trace”活动。
* 识别长时间运行的内核或计算密集型瓶颈。
* 右键单击时间线可快速启动配置文件。
* Nsight Compute 会自动过滤到选定的内核。

Nsight Systems 集成基本跟踪
Nsight Systems 集成基本跟踪

内联函数源表

Nsight Compute 针对内联函数提供了详细的源表分析功能。
* 可以按内联调用点或聚合到整个函数来分析度量指标。
* 使用 --inlineinfo 标志生成符号。
* 识别具体的性能不佳调用和异常值。
* 快速导航到内联调用点。

内联函数源表
内联函数源表

Hopper 线程块集群

Hopper 架构引入了线程块集群，以提供更高级别的局部性控制。
* 实现对大于单个 SM 上单个线程块粒度的可编程局部性控制。
* 集群使多个线程块能够跨多个 SM 并行运行，以同步和交换数据。
* 更多关于线程块集群的信息：
* [S51225] - CUDA: New Features and Beyond
* [S51119] - Optimizing Applications for Hopper Architecture

Hopper 线程块集群
Hopper 线程块集群

Nsight Compute 中的 Hopper 线程块集群

Nsight Compute 提供了对 Hopper 线程块集群的详细分析。

Nsight Compute 中的 Hopper 线程块集群
Nsight Compute 中的 Hopper 线程块集群

报告中的关键统计数据示例：
* 集群大小: 2.00 调度策略
* 网格大小: 140.00 每线程寄存器 [thread]
* 块大小: 384.00 静态共享内存每块 [byte/block]
* 线程 [thread]: 107,520.00 动态共享内存每块 [byte/block]
* 每 SM 波形数: 1.00 驱动程序共享内存每块 [byte/block]
* 函数缓存配置: cudaFuncCachePreferNone 共享内存配置大小 [byte]
* 理论占用率 [%]: 18.75 块限制寄存器 [block]
* 实现占用率 [%]: 12.00 块限制共享内存 [block]
* 实现占用率 [%]: 15.56 块限制 warp [block]
* 实现每 SM 活跃 warp 数 [warp]: 9.96 块限制 SM [block]
* 最大集群大小: 16.00 最大活跃集群
* 分支指令 [inst]: 1,570,420.00 分支效率 [%]
* 分支指令比率 [%]: 0.05 平均分歧分支

Hopper 张量内存加速器 (TMA)

Hopper 架构引入了张量内存加速器（TMA），以优化数据传输。
* 新的异步内存复制单元，用于将大量数据和多维张量从全局内存传输到共享内存，然后再传回。
* TMA 有助于为强大的新 H100 Tensor Cores 提供数据。
* TMA 释放线程来执行其他独立工作，而不是进行复制。
* 更多关于 TMA 的信息：
* [S51413] - Developing Optimal CUDA Kernels on Hopper Tensor Cores
* [S51119] - Optimizing Applications for Hopper Architecture

Hopper 张量内存加速器 (TMA)
Hopper 张量内存加速器 (TMA)

Nsight Compute 中对 Hopper TMA 的理解

Nsight Compute 能够帮助用户理解 Hopper TMA 的性能表现。

Nsight Compute 中对 Hopper TMA 的理解
Nsight Compute 中对 Hopper TMA 的理解

新增文档示例

为了降低性能优化的复杂性，Nsight Compute 提供了新的文档示例。
* 随 Nsight Compute 安装包提供。
* 包含可重现的说明、描述、预收集结果和源代码。
* 解释 CUDA 应用程序中常见的性能问题，例如“非合并全局访问和共享内存银行冲突”。

新增文档示例
新增文档示例

源页面指导 (即将推出)

即将推出的源页面指导功能将增强性能分析的便利性。
* 分支导航: 快速跳转到分支指令。
* 规则标记:
* 自动化分析和每行源代码的规则输出。
* 快速检测和导航到重要问题。
* 提供检测到的问题解释。

源页面指导
源页面指导

潜在加速计算 (即将推出)

即将推出的潜在加速计算功能将帮助用户更好地聚焦优化工作。

潜在加速计算
潜在加速计算

图示下方文本：
Issue Slot Utilization
每个调度器每周期能发出一个指令，但对于此内核，每个调度器平均每 16.5 个周期才发出一个指令。这可能导致硬件资源利用不足，并可能导致性能下降。在每个调度器最大 12 个 warp 中，此内核平均分配 8.28 个活跃 warp；但平均只有 0.09 个 warp 每周期符合条件。符合条件的 warp 是活跃 warp 中准备好发出下一条指令的子集。每个没有符合条件 warp 的周期都不会发出指令，并且发出槽位保持未使用。要增加符合条件 warp 的数量，可以通过检查 Warp 状态统计和源计数器部分中的主要停滞原因来减少活跃 warp 停滞的时间。
潜在加速: 16.48x

long_scoreboard
平均而言，此内核的每个 warp 花费 110.1 个周期等待 L1TEX (局部、全局、表面、纹理) 操作的记分牌依赖。找到产生被等待数据的指令以识别罪魁祸首。为了减少等待 L1TEX 数据访问的周期数，请验证内存访问模式是否针对目标架构进行了优化，尝试通过增加数据局部性（合并）来提高缓存命中率，或更改缓存配置。考虑将常用数据移动到共享内存。此停滞类型约占两次指令发出之间总平均 136.4 个周期的 80.7%。
潜在加速: 5.19x

CUPTI：CUDA 分析工具接口

CUPTI 更新

CUPTI (CUDA Profiling Tools Interface) 进行了多项更新，以提升其功能和性能。
* CUDA 设备启动的图支持。
* 支持 WSL2。
* 内存开销和性能改进。

CUPTI 更新
CUPTI 更新

CUPTI 配置示例：

void Profiler::Setup()
{
    // Enable CUPTI activities
    cuptiActivityEnable(CUPTI_ACTIVITY_KIND_DRIVER);
    cuptiActivityEnable(CUPTI_ACTIVITY_KIND_RUNTIME);
    cuptiActivityEnable(CUPTI_ACTIVITY_KIND_CONCURRENT_KERNELS);
    cuptiActivityEnable(CUPTI_ACTIVITY_KIND_MEMCPY);
    cuptiActivityEnable(CUPTI_ACTIVITY_KIND_MEMSET);

    // Register callbacks for buffer requested and completed
    cuptiActivityRegisterCallbacks(BufferRequested, BufferCompleted);
}

void Profiler::Flush()
{
    // Flush activity buffers
    cuptiActivityFlushAll();
}

从宏观到微观：全面整合 (From the Macro to the Micro: Putting it all together)

本节展示了一个从宏观（系统级）到微观（硬件级）逐步深入的性能分析方法，通过整合多种工具和视图来全面诊断性能瓶颈。

从宏观到微观
从宏观到微观

宏观视图：系统级通信分析

系统级通信分析
从宏观层面开始，分析MPI通信行为。图像展示了：
* 网络活动概览：热力图显示了不同网络通信器（NC-0, NC-1, NC-2, NC-3）在一段时间内的字节发送和接收情况，以及发送等待时间。
* MPI操作细节：具体到一个MPI_Allcast操作，显示其开始时间（0.104935s）、结束时间（0.165045s，持续109,210 µs）、线程ID（132434）、发送字节数（0）、接收字节数（4）、根节点（0）和通信域（MPI_COMM_WORLD）。

中观视图：CPU活动分析

CPU活动分析
在系统级通信分析的基础上，进一步深入到CPU层面的活动：
* CPU时间线视图：显示了多个CPU核心（如CPU 0, CPU 1, CPU 8, CPU 123）在不同时间点上的活动，包括CPU周期和指令执行。红线突出显示了特定CPU的繁忙或等待状态。
* 代码级性能分析：展示了与CPU活动相关的代码执行视图，可能用于识别热点函数或代码段。

微观视图：GPU活动分析

GPU活动分析
在CPU活动分析之后，进一步聚焦到GPU的性能：
* CUDA流与内核执行：展示了CUDA流（Stream 118）中内存操作和多个内核（如_Z6kernelv、launch_bounds、__cuda_memset）的执行情况。
* GPU SM指令与占用率：显示了SM指令、SM发放器、Tensor活跃时间和SM Warp占用率，这些指标有助于评估GPU计算单元的利用效率。
* GPU任务图：展示了GPU上不同计算和数据传输任务的依赖关系和执行流程。

细节视图：GPU工作负载与管道分析

GPU工作负载与管道分析
在GPU活动分析的基础上，提供更精细的GPU工作负载和管道分析：
* 内存工作负载分析：对GPU的内存资源进行详细分析，可能包括内存吞吐量、L2缓存命中率、L1TEX全局访问模式等，以识别内存访问瓶颈。
* 管道利用率：显示了计算单元（ALU, FMA）和内存单元的管道利用率，衡量GPU处理单元的效率。例如，ALU为25.0%，FMA为15.0%。
* 采样数据：提供指令（如133, 134, 135, 136, 137, 138行）的采样数据，显示其源模块、地址和具体指令内容，如BAKINC RBP、ISPYP_07 AND P1、ISPYP_07 AND P1等，用于识别微观层的性能瓶颈。

附加资源 (Additional Resources)

附加资源
附加资源

GTC开发者工具概览 (Developer Tools Across GTC)

GTC开发者工具概览
本节列举了在GTC（GPU Technology Conference）上与NVIDIA开发者工具相关的会议、实验和资源：

NVIDIA Logo
NVIDIA Logo